#include <stdio.h>
#include <stdlib.h>
#include <sys/types.h>
#include <sys/stat.h>
#include <fcntl.h>
#include <sys/mman.h>
#include <string.h>
#include <unistd.h>
#include <time.h>
#include <CL/cl.h>
#include <string>
#include <malloc.h>
#include "common.h"
#include "scanImpl.cpp"
#include "../../kernel/para.h"

#define BILLION 1000000000
#define	RATE	0.01

using namespace std;
extern const char * createProgram(string, int *);
struct statistic pp;

struct clContext context;


int main(int argc, char **argv){

	struct timespec start, afBuild, end;

	if(argc != 4){
		printf("./scanTest rate scanColNum whereColNum\n");
		exit(-1);
	}

	int rate = atoi(argv[1]) * RATE * 100;
	int scanColNum = atoi(argv[2]);
	int whereColNum = atoi(argv[3]);

	int * scanCol = (int *) malloc(sizeof(int) * scanColNum);
	int * whereCol = (int *) malloc(sizeof(int) * whereColNum);

	if(!scanCol || !whereCol){
		printf("Failed to allocate host memory\n");
		exit(-1);
	}

	char **scanTable = (char **) malloc(sizeof(char *) * scanColNum);
	char **whereTable = (char **) malloc(sizeof(char *) * whereColNum);

	if(!scanTable || !whereTable){
		printf("Failed to allocate host memory\n");
		exit(-1);
	}

	for(int i=0;i<scanColNum;i++){
		char buf[32] = {0};
		sprintf(buf,"scanTable%d",i);
		scanCol[i] = open(buf,O_RDONLY);

		int tmpSize = lseek(scanCol[i],0,SEEK_END); 
		scanTable[i] = (char*) memalign(128,tmpSize);
		if(!scanTable[i]){
			printf("Failed to allocate host memory for scanTable\n");
			exit(-1);
		}
		char *tmp = (char *)mmap(0,tmpSize,PROT_READ,MAP_SHARED,scanCol[i],0);
		memcpy(scanTable[i], tmp, tmpSize);
		munmap(tmp,tmpSize);
	}

	for(int i=0;i<whereColNum;i++){
		char buf[32] = {0};
		sprintf(buf,"whereTable%d",i);
		whereCol[i] = open(buf,O_RDONLY);

		int tmpSize = lseek(whereCol[i],0,SEEK_END); 
		whereTable[i] = (char*) memalign(128,tmpSize);
		if(!whereTable[i]){
			printf("Failed to allocate host memory for whereTable\n");
			exit(-1);
		}
		char *tmp = (char *)mmap(0,tmpSize,PROT_READ,MAP_SHARED,whereCol[i],0);
		memcpy(whereTable[i], tmp, tmpSize);
		munmap(tmp,tmpSize);
	}

	int whereSize = lseek(whereCol[0],0,SEEK_END);
	int totalTupleNum = whereSize / sizeof(int); 

	int psc = 0;
	const char *ps = createProgram("../kernel/select.cl",&psc);

	clock_gettime(CLOCK_REALTIME,&start);

	size_t localSize = LOCAL;
	size_t globalSize = GLOBAL*localSize;
	size_t threadNum = globalSize;
	cl_int error = 0;
        cl_device_id device;

        cl_uint numP;
        clGetPlatformIDs(0,NULL,&numP);

        cl_platform_id * pid = new cl_platform_id[numP];
        clGetPlatformIDs(numP, pid, NULL);

        clGetDeviceIDs(pid[PID],DEVICETYPE,1,&device,NULL);

        context.context = clCreateContext(0,1,&device,NULL,NULL,&error);
        if(error != CL_SUCCESS){
                printf("Error creating context\n");
                exit(-1);
        }

        context.queue = clCreateCommandQueue(context.context, device, 0, &error);
        if(error != CL_SUCCESS){
                printf("Error creating command queue\n");
                exit(-1);
        }

        context.program = clCreateProgramWithSource(context.context, psc, (const char **)&ps, 0, &error);
        if(error != CL_SUCCESS){
                printf("Error creating program\n");
                exit(-1);
        }

	error = clBuildProgram(context.program, 0, 0 , 0 , 0, 0);
        if(error != CL_SUCCESS){
                printf("Error building program\n");
                exit(-1);
        }


	cl_mem gpuFilter = clCreateBuffer(context.context, CL_MEM_READ_WRITE, sizeof(int)*totalTupleNum, NULL, &error);
	cl_mem gpuPsum = clCreateBuffer(context.context, CL_MEM_READ_WRITE, sizeof(int)*threadNum, NULL, &error);
	cl_mem gpuCount = clCreateBuffer(context.context, CL_MEM_READ_WRITE, sizeof(int)*threadNum, NULL, &error);
	cl_mem result = clCreateBuffer(context.context, CL_MEM_READ_WRITE, totalTupleNum * sizeof(int), NULL, &error);

	int count;

	clock_gettime(CLOCK_REALTIME,&afBuild);

	cl_mem column = clCreateBuffer(context.context, CL_MEM_READ_ONLY|CL_MEM_USE_HOST_PTR, whereSize, whereTable[0], &error);
	//clEnqueueWriteBuffer(context.queue,column,CL_TRUE,0,whereSize,whereTable[0],0,0,0);
	context.kernel = clCreateKernel(context.program, "genScanFilter_init", 0);

	clSetKernelArg(context.kernel, 0, sizeof(cl_mem), (void *)&column);
	clSetKernelArg(context.kernel, 1, sizeof(int), (void *)&totalTupleNum);
	clSetKernelArg(context.kernel, 2, sizeof(int), (void *)&rate);
	clSetKernelArg(context.kernel, 3, sizeof(cl_mem), (void *)&gpuFilter);
	clEnqueueNDRangeKernel(context.queue,context.kernel, 1,0, (size_t *)&globalSize,&localSize,0,0,0);


	for(int i=1;i<whereColNum;i++){
		column = clCreateBuffer(context.context, CL_MEM_READ_ONLY|CL_MEM_USE_HOST_PTR, whereSize, whereTable[0], &error);
		//clEnqueueWriteBuffer(context.queue,column,CL_TRUE,0,whereSize,whereTable[1],0,0,0);
		context.kernel = clCreateKernel(context.program, "genScanFilter_and", 0);

		clSetKernelArg(context.kernel, 0, sizeof(cl_mem), (void *)&column);
		clSetKernelArg(context.kernel, 1, sizeof(int), (void *)&totalTupleNum);
		clSetKernelArg(context.kernel, 2, sizeof(int), (void *)&rate);
		clSetKernelArg(context.kernel, 3, sizeof(cl_mem), (void *)&gpuFilter);
		clEnqueueNDRangeKernel(context.queue,context.kernel, 1,0, (size_t *)&globalSize,&localSize,0,0,0);
	}

	context.kernel = clCreateKernel(context.program, "countScanNum", 0);
	clSetKernelArg(context.kernel, 0, sizeof(cl_mem), (void *)&gpuFilter);
	clSetKernelArg(context.kernel, 1, sizeof(int), (void *)&totalTupleNum);
	clSetKernelArg(context.kernel, 2, sizeof(cl_mem), (void *)&gpuCount);
	clEnqueueNDRangeKernel(context.queue,context.kernel, 1,0, (size_t *)&globalSize,&localSize,0,0,0);

	scanImpl(gpuCount,threadNum,gpuPsum,&context, &pp);

	int tmp1, tmp2;
        clEnqueueReadBuffer(context.queue, gpuCount, CL_TRUE, sizeof(int)*(threadNum-1), sizeof(int), &tmp1,0,0,0);
        clEnqueueReadBuffer(context.queue, gpuPsum, CL_TRUE, sizeof(int)*(threadNum-1), sizeof(int), &tmp2,0,0,0);

        count = tmp1+tmp2;

	printf("%d\n",count);

	for(int i=0;i<scanColNum;i++){
		column = clCreateBuffer(context.context, CL_MEM_READ_ONLY, totalTupleNum * sizeof(int), scanTable[i], &error);
		//clEnqueueWriteBuffer(context.queue,column,CL_TRUE,0,whereSize,scanTable[i],0,0,0);


		context.kernel = clCreateKernel(context.program, "scan", 0);
		clSetKernelArg(context.kernel, 0, sizeof(cl_mem), (void *)&column);
		clSetKernelArg(context.kernel, 1, sizeof(int), (void *)&totalTupleNum);
		clSetKernelArg(context.kernel, 2, sizeof(cl_mem), (void *)&gpuPsum);
		clSetKernelArg(context.kernel, 3, sizeof(int), (void *)&count);
		clSetKernelArg(context.kernel, 4, sizeof(cl_mem), (void *)&gpuFilter);
		clSetKernelArg(context.kernel, 5, sizeof(cl_mem), (void *)&result);
		clEnqueueNDRangeKernel(context.queue,context.kernel, 1,0, (size_t *)&globalSize,&localSize,0,0,0);

		clEnqueueMapBuffer(context.queue,result,CL_TRUE,CL_MAP_READ,0,sizeof(int)*count, 0,0,0,0);
	}

	clFinish(context.queue);
	clReleaseCommandQueue(context.queue);

	clock_gettime(CLOCK_REALTIME,&end);
        double timeE = (end.tv_sec -  start.tv_sec)* BILLION + end.tv_nsec - start.tv_nsec;
        printf("Table scan CPU time: %lf\n", timeE);

        timeE = (end.tv_sec -  afBuild.tv_sec)* BILLION + end.tv_nsec - afBuild.tv_nsec;
        printf("Table scan CPU time without build: %lf\n", timeE);


	for(int i=0;i<scanColNum;i++){
		close(scanCol[i]);
		free(scanTable[i]);
	}

	for(int i=0;i<whereColNum;i++){
		close(whereCol[i]);
		free(whereTable[i]);
	}

	free(scanTable);
	free(whereTable);

}

